feat: Triton fused SoftSignGLU kernel for LYNXNet2[WIP]#312
feat: Triton fused SoftSignGLU kernel for LYNXNet2[WIP]#312KakaruHayate wants to merge 16 commits into
Conversation
Fuse nn.Linear + SoftSignGLU into single Triton kernel launches, reducing HBM traffic by ~5 GB/step on LYNXNet2 acoustic backbone. Supported architecture: LYNXNet2 with SoftSignGLU activation only. Key components: - modules/kernels/fused_linear_softsign_glu.py: Fused forward/backward kernels with weight-split strategy to avoid cross-program communication. SoftSignGLU is numerically exact (no Taylor approximation needed). - modules/kernels/integration.py: Non-invasive monkey-patch with automatic eval-mode fallback for ONNX export. - Autotune warmup at training start to avoid first-step compilation lag. Usage: Set 'use_fused_kernels: true' in config (default: true). ONNX export works automatically via model.eval() fallback.
|
@codex review |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: d8b261c697
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
- Variance model backbones (K=384/512) are too small for Triton fusion to benefit; skip patching to avoid compilation overhead. - Warmup with M=50000 (matching max_batch_frames) so Triton cache hits on the first real training step instead of recompiling. - Reduce autotune configs from 6+5+4=15 to 3+3+2=8, cutting compile time by ~50%.
Warmup at init time runs before Lightning moves the model to GPU, so Triton cannot compile kernels there. Remove the calls entirely; compilation happens naturally on the first real training step. Also fix default glu_type to softsign_glu in both tasks.
Warmup at init time runs before Lightning moves the model to GPU, so Triton cannot compile kernels there. Remove the calls entirely; compilation happens naturally on the first real training step. Also fix default glu_type to softsign_glu in both tasks.
- P1: keep use_fused_kernels opt-in (false default) for backward compat - P1: split weight at midpoint not at K, supporting expansion_factor != 1 - P2: add N param (GLU output dim) separate from K (input/conraction dim) - P2: save ctx.N for backward wrapper calculations
|
@codex review |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 9256c0b8dd
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
- Remove extra N arg in _softsign_glu_bwd_elem_kernel call (M,N,N → M,N) - Read variance glu_type from nested predictor config paths - Assert both predictors use softsign_glu before patching
|
@autumn-2-net hello? |
|
@codex review |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 9c17e7b28f
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
9c17e7b to
ade184c
Compare
… fewer autotune configs, realistic warmup M
- weight split at midpoint (support expansion_factor != 1) - variance nested glu_type from predictor configs - remove useless warmup from __init__ (model on CPU, no compilation) - use_fused_kernels default to false
|
@autumn-2-net hello |
- Backward: remove custom Triton GEMM (3.4x slower than cuBLAS), use cuBLAS matmuls with out= preallocation for grad_x/grad_w - Autotune key: bucket M by next_power_of_2 to prevent ~2.6s per-batch re-benchmark under variable frame counts (DsBatchSampler) - Forward configs: add large tiles for Ada/Blackwell (4090/5090) + GROUP_M swizzle for L2 reuse; small tiles retained for Turing - Elem backward kernel: fix key=['M','K'] -> key=['N'], fix param name mismatch (M,K -> M,N at call site) - bf16 fallback: tl.dot bf16 crashes on pre-Ampere (Turing sm_75); detect at runtime, fall back to eager path - Warmup: sweep M buckets under torch.autocast, derive cond hidden size from backbone instead of hardcoding 384 - Integration: restrict to softsign_glu only; glu_type default 'swiglu'; per-predictor patching for variance (no blanket assertion) - Acoustic/variance tasks: on_fit_start warmup hook; glu_type fallback Benchmark (RTX 2070, fp16, K=1024): Before: 0.69x (fwd+bwd), 2669ms on new M (135x slower) After: 1.42x (fwd+bwd), 15ms on new M (1.41x faster)
|
@autumn-2-net 人呢? |
- Remove global TF32 side-effects from acoustic_task / variance_task (allow_tf32 / set_float32_matmul_precision were unrelated to this PR and silently change fp32 training for every user). - Revert 6 config glu_type from softsign_glu back to atanglu. softsign_glu stays available as an explicit opt-in; switching the default is an architectural decision that needs A/B evidence first. - Add use_fused_kernels_variance: false to base.yaml so the key is discoverable instead of buried in a code comment. - Protect kernel import on Windows / no-Triton: try/import triton in fused_linear_softsign_glu.py + capability gate for sm < 70. Guard Fn class + bwd kernel inside if _TRITON_AVAILABLE so the module loads cleanly when Triton is absent. - Handle bias=None in Fn.forward (would AttributeError on .split()). - Remove no-op view in forward(x.dim()>2): left/gate are already [M,N]. - Add one-time-per-(dtype,K) warning when fused falls back to eager (prevents fp32 / unsupported-dtype confusion). - Integration: replace net closure with self.net[] to avoid the deepcopy stale-closure issue (EMA / SWA safety). - Warmup: run no_grad forward only (DDP-safe); drop unused params (glu_type, num_channels) from signature. - Task files: wrap patch calls in try/except ImportError so a missing Triton install logs a clear message and continues in eager mode; log autocast_dtype=None case so fp32 users know fused is inactive. - Test: add assert thresholds so numerical regressions fail loudly.
- VarianceTask: add missing on_fit_start warmup (mirrors AcousticTask, sweeps pitch_predictor and variance_predictor over denoise_fn / velocity_fn). Previously use_fused_kernels_variance patched blocks but never warmed the autotune cache. - _test(): compare fused AND eager against a shared fp32 reference instead of fused-vs-eager fp16. fp16 eager itself deviates ~1e-2 relative from fp32 at K=512 (different accumulation order), so the old absolute threshold conflated rounding with kernel bugs. New criterion: fused error <= 2x eager error. Verified on RTX 2070 (sm_75, torch 2.12.1+cu130, triton 3.2.0): - kernel test: fused error consistently BELOW eager fp16 error (fp32 accumulators), K=256/512/1024 fwd+bwd all pass - integration test: block fwd diff 1.95e-3, grad diff 2.20e-3 - eval-mode fallback matches train-mode fused output - deepcopy safety: copied block uses own weights (stale-closure fixed) - fallbacks: bias=None, CPU, fp32 (one-time warning), bf16-on-Turing - simulated no-Triton import: module loads, eager fallback works - warmup: no_grad forward-only, no grads created; full autocast fp16 training step through patched LYNXNet2 works
If the modules.kernels.integration import fails, rank_zero_info (imported on the following line inside the same try) is also unbound; import it locally in the except clause like variance_task does.
…ory trick Eager (common_layers.py): - SoftSignGLUFunction: ATanGLUFunction-style custom backward. softsign'(x) = (1-|softsign(x)|)^2, so both partials are precomputable in forward; backward is two pure multiplies, saves one tensor vs naive autograd. SoftSignGLU now uses it in training mode. - DoubleSoftSignGLU: y = softsign(out) * softsign(gate). softsign applied to the whole Linear output then split (elementwise => equivalent). Custom Function precomputes both partials into one [.., 2N] buffer. - lynxnet2: register glu_type 'double_softsign_glu'. Kernel (fused_linear_softsign_glu.py): - IS_DOUBLE constexpr on fwd + bwd-elem kernels; Triton compiles two specializations, zero overhead in the single-gate path. - fused_linear_softsign_glu(x, w, b, is_double=False) public API; eager fallback handles both modes. - integration.py: _FUSABLE_GLU_TYPES = (softsign_glu, double_softsign_glu). Verified on RTX 2070: - eager Functions vs naive autograd in fp64: max diff < 2e-15 (both) - double-mode kernel vs fp32 reference: fused error 2-5x SMALLER than eager fp16 (fwd 3.2e-4 vs 1.8e-3) — fp32 accumulators - single-gate regression: unchanged, all pass - block integration both modes: fwd/grad diff ~2e-3, eval fallback exact
One switch for both acoustic and variance training. The original
rationale for a separate variance key ('backbones too small to benefit')
did not hold up in benchmarks: at 24k frames/batch the variance backbone
gains 1.35x (single) / 1.62x (double softsign). The per-predictor
glu_type check already prevents patching non-softsign backbones, so the
extra key only added confusion. base.yaml comment now documents the
actual trade-off (benefit scales with max_batch_frames).
Fuse nn.Linear + SoftSignGLU into single Triton kernel launches, reducing HBM traffic by ~5 GB/step on LYNXNet2 acoustic backbone.
Supported architecture: LYNXNet2 with SoftSignGLU activation only.
Key components:
Usage:
Set 'use_fused_kernels: true' in config (default: true). ONNX export works automatically via model.eval() fallback.